`cuda::is_trivially_copyable` by fbusato · Pull Request #8265 · NVIDIA/cccl

fbusato · 2026-04-01T19:51:06Z

Description

Followup of the discussion in

Recover warp_shuffle original behavior (revert #8210) #8254

The PR introduces cuda::is_trivially_copyable_relaxed to support types that are actually trivially copyable but not recognized by the C++ std::is_trivially_copyable type trait.

The new trait supports:

trivially copyable types.
extended floating-point types.
extended floating-point vector types.
Raw arrays, cuda::std::array.
cuda::std::pair.
cuda::std::tuple.
Nested compositions of the above cases.

Potentially affected paths by `std::is_trivially_copyable`

unsupported cases:

cuda/std/__atomic/types/base.h:39
cuda/std/__atomic/types/reference.h:39
cuda/__memcpy_async/memcpy_async_barrier.h:61
cuda/__memcpy_async/memcpy_async_tx.h:58
cuda/__container/buffer.h:111
cuda/__algorithm/copy.h:73
cuda/__algorithm/fill.h:42
cuda/std/string_view:147

fallback to slower path:

cuda/std/__algorithm/copy.h:102
cuda/std/__algorithm/copy_backward.h:50
cuda/std/__algorithm/move.h:54
cuda/std/__algorithm/move_backward.h:53
cuda/std/__bit/bit_cast.h:55
cub/device/dispatch/kernels/kernel_histogram.cuh:81
cub/detail/uninitialized_copy.cuh:33
thrust/type_traits/is_trivially_relocatable.h:213
cudax/include/cuda/experimental/__kernel/kernel_ref.cuh:54

miscco

I am not too excited about this.

We need to be really careful here because the compiler may also break in some of those cases.

I did some experiments in #3183

bernhardmgruber · 2026-04-01T20:38:11Z

I wanted to start a discussion for some time to add such a trait, but I would really keep this feature internally for now until it is sufficiently baked. I would really like this feature to make thrust::is_trivially_relocatable and THRUST_PROCLAIM_TRIVIALLY_RELOCATABLE obsolete, because those two are often used wrongly, so we should make sure the new trait here can replace it.

davebayer

I am not a fan of this trait. We bend C++ rules to fix poorly designed nvfp types. We would have to basically change every use of is_trivially_copyable and is_trivially_copy_constructible to this trait to make it work consistently in libcu++.

I don't think this is a good idea, we should rather insist of the nvfp types being fixed.

docs/libcudacxx/extended_api/type_traits/is_trivially_copyable_relaxed.rst

davebayer · 2026-04-01T21:12:53Z

docs/libcudacxx/extended_api/type_traits/is_trivially_copyable_relaxed.rst

+Users may specialize ``cuda::is_trivially_copyable_relaxed`` for their own types whose memory representation is safe to copy
+with ``memcpy`` but that the compiler does not consider trivially copyable.


Isn't this just UB?

not for the types that we care about. Said that, the user could provide an object that triggers UB. I can highlight this point in the documentation but we cannot do anything to explicitly prevent it.

fbusato · 2026-04-01T22:26:29Z

We need to be really careful here because the compiler may also break in some of those cases.

I did some experiments in #3183

Integrated and extended the tests that you point out. Everything works

fbusato · 2026-04-01T22:27:40Z

I wanted to start a discussion for some time to add such a trait, but I would really keep this feature internally for now until it is sufficiently baked. I would really like this feature to make thrust::is_trivially_relocatable and THRUST_PROCLAIM_TRIVIALLY_RELOCATABLE obsolete, because those two are often used wrongly, so we should make sure the new trait here can replace it.

the actual issue is that the user can extend the trait to custom types. It needs to be public

fbusato · 2026-04-01T22:29:56Z

I am not a fan of this trait. We bend C++ rules to fix poorly designed nvfp types. We would have to basically change every use of is_trivially_copyable and is_trivially_copy_constructible to this trait to make it work consistently in libcu++.
I don't think this is a good idea, we should rather insist of the nvfp types being fixed.

there are no plan for that. nvfp types are not trivially copyable for optimization purposes.

jrhemstad · 2026-04-01T23:03:39Z

we should rather insist of the nvfp types being fixed.

This isn't going to happen and we shouldn't delude ourselves into thinking it ever will.

docs/libcudacxx/extended_api/type_traits/is_trivially_copyable_relaxed.rst

libcudacxx/include/cuda/__type_traits/is_trivially_copyable_relaxed.h

libcudacxx/include/cuda/__type_traits/is_trivially_copyable.h

…_relaxed.rst Co-authored-by: Giannis Gonidelis <gonidelis@hotmail.com>

…laxed-type-traits

libcudacxx/test/libcudacxx/cuda/type_traits/is_trivially_copyable.aggr.pass.cpp

docs/libcudacxx/extended_api/type_traits/is_trivially_copyable.rst

fbusato · 2026-04-08T19:55:17Z

libcudacxx/include/cuda/std/__bit/bit_cast.h

+[[nodiscard]] _CCCL_API inline _CCCL_CONSTEXPR_BIT_CAST _To bit_cast(const _From& __from) noexcept
+{
+#if defined(_CCCL_BUILTIN_BIT_CAST)
+  if constexpr (::cuda::std::is_trivially_copyable_v<_To> && ::cuda::std::is_trivially_copyable_v<_From>)


here the trick. We know that __nv_bfloat2 and __half2 are not trivially copyable but, we also know, that these types are not available in constant expressions so we can dispatch depending on this property

github-actions · 2026-04-09T02:06:25Z

😬 CI Workflow Results

🟥 Finished in 4h 08m: Pass: 87%/99 | Total: 5d 00h | Max: 4h 07m | Hits: 37%/249239

See results here.

miscco · 2026-04-09T08:00:02Z

libcudacxx/include/cuda/__type_traits/is_trivially_copyable.h

@@ -0,0 +1,84 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.


Should be the libcu++ license

miscco · 2026-04-09T08:01:44Z

libcudacxx/include/cuda/__type_traits/is_trivially_copyable.h

+_CCCL_BEGIN_NAMESPACE_CUDA
+
+template <typename _Tp, typename = void>
+constexpr bool __is_aggregate_trivially_copyable_v = false;


Critical: All those need to be inline constexpr

miscco · 2026-04-09T08:03:36Z

libcudacxx/include/cuda/std/__bit/bit_cast.h

+#if defined(_CCCL_BUILTIN_BIT_CAST)
+  if constexpr (::cuda::std::is_trivially_copyable_v<_To> && ::cuda::std::is_trivially_copyable_v<_From>)
+  {
+    return _CCCL_BUILTIN_BIT_CAST(_To, __from);
+  }
+  else
+  {
+    return ::cuda::std::__bit_cast_memcpy<_To>(__from);
+  }
+#else // ^^^ _CCCL_BUILTIN_BIT_CAST ^^^ / vvv !_CCCL_BUILTIN_BIT_CAST vvv
+  return ::cuda::std::__bit_cast_memcpy<_To>(__from);
 #endif // !_CCCL_BUILTIN_BIT_CAST


We can shorten this

Suggested change

#if defined(_CCCL_BUILTIN_BIT_CAST)

if constexpr (::cuda::std::is_trivially_copyable_v<_To> && ::cuda::std::is_trivially_copyable_v<_From>)

{

return _CCCL_BUILTIN_BIT_CAST(_To, __from);

}

else

{

return ::cuda::std::__bit_cast_memcpy<_To>(__from);

}

#else // ^^^ _CCCL_BUILTIN_BIT_CAST ^^^ / vvv !_CCCL_BUILTIN_BIT_CAST vvv

return ::cuda::std::__bit_cast_memcpy<_To>(__from);

#endif // !_CCCL_BUILTIN_BIT_CAST

#if defined(_CCCL_BUILTIN_BIT_CAST)

if constexpr (::cuda::std::is_trivially_copyable_v<_To> && ::cuda::std::is_trivially_copyable_v<_From>)

{

return _CCCL_BUILTIN_BIT_CAST(_To, __from);

}

else

#endif // _CCCL_BUILTIN_BIT_CAST

{

return ::cuda::std::__bit_cast_memcpy<_To>(__from);

}

miscco · 2026-04-09T08:04:53Z

libcudacxx/include/cuda/std/__type_traits/aggregate_members.h

+#if defined(_CCCL_BUILTIN_STRUCTURED_BINDING_SIZE)
+
+// Returns the number of aggregate members, or `-1` if the type is not an aggregate.
+template <typename _Tp, ::cuda::std::enable_if_t<_CCCL_BUILTIN_STRUCTURED_BINDING_SIZE(_Tp) >= 0, int> = 0>


Critical: This is not defined in libcu++ but only in cudax currently.

We need to move the builtin definition into libcu++

miscco · 2026-04-09T08:05:10Z

libcudacxx/include/cuda/std/__type_traits/aggregate_members.h

+
+// Returns the number of aggregate members, or `-1` if the type is not an aggregate.
+template <typename _Tp, ::cuda::std::enable_if_t<_CCCL_BUILTIN_STRUCTURED_BINDING_SIZE(_Tp) >= 0, int> = 0>
+constexpr int __aggregate_arity_v = _CCCL_BUILTIN_STRUCTURED_BINDING_SIZE(_Tp);


Critical: missing inline

Suggested change

constexpr int __aggregate_arity_v = _CCCL_BUILTIN_STRUCTURED_BINDING_SIZE(_Tp);

inline constexpr int __aggregate_arity_v = _CCCL_BUILTIN_STRUCTURED_BINDING_SIZE(_Tp);

miscco · 2026-04-09T08:05:34Z

libcudacxx/include/cuda/std/__type_traits/aggregate_members.h

+
+// Returns the number of aggregate members, or `-1` if the type is not an aggregate.
+template <typename _Tp>
+constexpr int __aggregate_arity_v = int{sizeof(*__aggregate_arity_impl<_Tp>{}())} - 2;


Critical: missing inline

Suggested change

constexpr int __aggregate_arity_v = int{sizeof(*__aggregate_arity_impl<_Tp>{}())} - 2;

inline constexpr int __aggregate_arity_v = int{sizeof(*__aggregate_arity_impl<_Tp>{}())} - 2;

miscco · 2026-04-09T08:08:18Z

libcudacxx/include/cuda/std/__type_traits/aggregate_members.h

Please move this out of the PR.

This is a nontrivial piece of code and the current implementation seems really costly.

I want this in a separate PR so that we can properly review it

miscco · 2026-04-09T08:09:18Z

libcudacxx/test/libcudacxx/cuda/type_traits/is_trivially_copyable.aggr.pass.cpp

+  static_assert(cuda::is_trivially_copyable<const _Tp>::value);
+  static_assert(cuda::is_trivially_copyable_v<_Tp>);
+  static_assert(cuda::is_trivially_copyable_v<const _Tp>);
+}


Important: This should test that we can perform a bit_cast from this type to a similarly sized type

miscco · 2026-04-09T08:10:31Z

libcudacxx/test/libcudacxx/cuda/type_traits/is_trivially_copyable.pass.cpp

+  static_assert(cuda::is_trivially_copyable<T>::value);
+  static_assert(cuda::is_trivially_copyable<const T>::value);
+  static_assert(cuda::is_trivially_copyable_v<T>);
+  static_assert(cuda::is_trivially_copyable_v<const T>);


Missing tests for bitcast back and forth

is_trivially_copyable_relaxed

ea3f956

fbusato self-assigned this Apr 1, 2026

fbusato added this to CCCL Apr 1, 2026

fbusato added the libcu++ For all items related to libcu++ label Apr 1, 2026

fbusato requested a review from a team as a code owner April 1, 2026 19:51

github-project-automation bot moved this to Todo in CCCL Apr 1, 2026

fbusato requested a review from wmaxey April 1, 2026 19:51

cccl-authenticator-app bot moved this from Todo to In Review in CCCL Apr 1, 2026

add documentation

9dd60db

fbusato requested a review from a team as a code owner April 1, 2026 19:55

fbusato requested a review from gonidelis April 1, 2026 19:55

fbusato added 2 commits April 1, 2026 13:02

describe custom specialization

179a81b

move to internal function

2c33c2e

miscco requested changes Apr 1, 2026

View reviewed changes

github-project-automation bot moved this from In Review to In Progress in CCCL Apr 1, 2026

davebayer requested changes Apr 1, 2026

View reviewed changes

address padding

fbade34

fbusato added 2 commits April 1, 2026 15:33

fix clang

deb622c

do not handle volatile

a553bb6

fbusato added 2 commits April 1, 2026 16:19

unused-local-typedef

de42a45

fix clang pragma

4e7873d

fbusato moved this from In Progress to In Review in CCCL Apr 1, 2026

This comment has been minimized.

Sign in to view

fbusato added 2 commits April 2, 2026 10:36

simplify conditions

6e5021f

improve documentation

8ff54f4